{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "Q6RWcI4z-2pW"
   },
   "source": [
    "# CUDA C/C++을 이용한 가속화 애플리케이션 \n",
    "\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "qe_dKUYg8vJ5"
   },
   "source": [
    "![CUDA](./images/CUDA_Logo.jpg)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "P4gex5U_8qbH"
   },
   "source": [
    "\n",
    "\n",
    "가속화 컴퓨팅은 CPU만을 사용한 컴퓨팅을 대체하는 새로운 베스트 프랙티스입니다. 가속화 컴퓨팅을 통해 이룬 수많은 개선 사례, 증가일로에 있는 가속화 애플리케이션에 대한 요청, 프로그래밍을 쉽게 할 수 있도록 돕는 각종 규약, 그리고 이를 지원하는 하드웨어의 지속적인 발전 등이 변화를 이끌고 있습니다.\n",
    "\n",
    "뛰어난 성능과 사용 편의성 관점에서 가속화 컴퓨팅을 바라본다면 그 중심에는 [CUDA](https://developer.nvidia.com/about-cuda) 컴퓨팅 플랫폼이 자리 잡고 있습니다. CUDA는 C, C++, 파이썬, 포트란과 같은 기존의 언어를 확장하여, 가속화되고 고도로 병렬화된 코드를 세계에서 가장 효과적인 병렬 프로세서인 NVIDIA GPU 상에서 실행할 수 있도록 만들어 주는 코딩 패러다임입니다. CUDA는 적은 노력만으로도 애플리케이션을 극적으로 가속화할 수 있으며, [DNN](https://developer.nvidia.com/cudnn), [BLAS](https://developer.nvidia.com/cublas), [그래프 애널리틱스](https://developer.nvidia.com/nvgraph), [FFT](https://developer.nvidia.com/cufft) 등 다양한 최적화 라이브러리로 구성된 생태계를 지니고 있습니다. 또한 강력한 [커맨드라인](http://docs.nvidia.com/cuda/profiler-users-guide/index.html#nvprof-overview) 및 [비쥬얼 프로파일러](http://docs.nvidia.com/cuda/profiler-users-guide/index.html#visual)를 제공합니다. CUDA는 또한 [전산유체역학](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=10,12,16,17,19,51,53,71,87,121,124,156,157,195,202,203,204,312,339,340,395,407,448,4&search=), [분자역학](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398&search=), [양자화학](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398&search=), [물리학](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216&search=), 고성능 컴퓨팅 등의 분야를 위한 [세계에서 가장 우수한 애플리케이션](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=58,59,60,293,98,172,223,227,228,265,487,488,114,389,220,258,461&search=)을 지원합니다.\n",
    "\n",
    "CUDA를 학습함으로써 여러분도 자신의 애플리케이션을 가속화할 수 있습니다. 가속화된 애플리케이션은 CPU만을 사용하여 실행할 때보다 훨씬 빠르게 작동하며, CPU의 성능 제약으로 인해 할 수 없었던 계산을 가능하게 합니다. 본 강좌에서 여러분은 CUDA C/C++를 이용하여 가속화 애플리케이션을 프로그래밍하는 방법을 배우고, CPU 기반의 애플리케이션을 가속화하여 성능 이득을 얻는 연습을 통해 컴퓨팅의 새로운 영역으로 첫발을 내딛게 될것입니다. \n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "mf2PZ8ds82wG"
   },
   "source": [
    "---\n",
    "\n",
    "## 선행지식 \n",
    "\n",
    "강좌의 이해를 위하여 여러분은 다음 사항을 알고 있어야 합니다.\n",
    "\n",
    "* C 언어로 변수를 선언하고 반복문을 만들고 if/else 구문을 사용하기\n",
    "* C 언어로 함수를 정의하고 호출하기\n",
    "* C 언어로 배열을 할당하기\n",
    "\n",
    "CUDA와 관련된 선행지식은 필요하지 않습니다.\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "WbF1om_M85-f"
   },
   "source": [
    "---\n",
    "\n",
    "## 학습목표\n",
    "\n",
    "강좌를 마치면 여러분은 다음 사항을 수행할 수 있게 될 것입니다.\n",
    "\n",
    "* CPU 함수를 호출하고 GPU **커널(kernel)을 구동(launch)**하는 C/C++ 프로그램을 작성, 컴파일, 실행하기\n",
    "* **실행 설정(execution configuration)**을 통해 병렬 **스레드 계층(thread hierarchy)** 제어하기 \n",
    "* 직렬로 실행되는 반복문을 GPU 상에서 병렬로 실행되도록 수정하기\n",
    "* CPU와 GPU를 위한 메모리를 할당하고 해제하기\n",
    "* CUDA 코드에서 생성된 오류 처리하기\n",
    "* CPU-only 애플리케이션 가속화하기\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "flnEclNB86IQ"
   },
   "source": [
    "---\n",
    "\n",
    "## 가속화 시스템\n",
    "\n",
    "가속화 시스템은 이종(heterogeneous) 시스템이라고 부르기도 합니다. 이는 시스템이 서로 이질적인 CPU와 GPU의 합으로 구성되기 때문에 생긴 이름입니다. 가속화 시스템은 CPU 프로그램을 수행하고, 이는 다시 GPU가 제공하는 고도의 병렬화 기능을 활용하기 위한 함수들을 구동하는 구조입니다. 본 강좌의 개발환경은 NVIDIA GPU를 포함한 가속화 시스템입니다. GPU에 대한 정보는 `nvidia-smi` 커맨드라인 명령어로 확인할 수 있습니다. 여기에서 smi는 systems management interface의 약자입니다. 아래 코드 실행 셀에서 `nvidia-smi` 명령을 입력하고 `CTRL+ENTER`를 눌러서 실행하세요. 본 강좌를 진행하면서 코드를 실행해야 할 때마다 이 셀들을 보게 될 것입니다. 실행된 결과의 출력은 코드가 실행된 후 셀의 아랫부분에 표시될 것입니다. 아래 코드를 실행한 후 출력에서 GPU의 이름을 잘 찾아서 확인해 보세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "f0gnvlRd-2pX"
   },
   "outputs": [],
   "source": [
    "!nvidia-smi"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "0RR5M8gi-2pa"
   },
   "source": [
    "---\n",
    "\n",
    "## GPU 가속화 시스템 vs. CPU-only 시스템\n",
    "\n",
    "다음 슬라이드는 앞으로 공부할 내용을 개략적인 수준에서 시각적으로 보여줍니다. 주제를 보다 상세히 다루기에 앞서 슬라이드를 클릭하시면서 살펴 보세요. \n",
    "\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {
     "base_uri": "https://localhost:8080/",
     "height": 591
    },
    "colab_type": "code",
    "executionInfo": {
     "elapsed": 1041,
     "status": "ok",
     "timestamp": 1541021727312,
     "user": {
      "displayName": "Hyungon Ryu KR",
      "photoUrl": "https://lh5.googleusercontent.com/-wtRHVbivPO4/AAAAAAAAAAI/AAAAAAAAABk/eM5PPjT2Row/s64/photo.jpg",
      "userId": "08585482373097815411"
     },
     "user_tz": -540
    },
    "id": "cGpgyw5X-2pc",
    "outputId": "4d0e0507-1feb-4737-b336-c6a20eda7677"
   },
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/e/2PACX-1vTdUbQjoEYAtcPCMX4ZVLa9gE0rbO3ODClJqgtzRaXoFgVmTJrOkXGDAYmA0BsuTEaokTASv84JsKLm/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "6M1ROEpR-2pf"
   },
   "source": [
    "---\n",
    "\n",
    "## GPU를 위한 애플리케이션 코드 작성하기\n",
    "\n",
    "CUDA는 다양한 범용 프로그래밍 언어 확장을 지원합니다. 본 강좌에서는 그 중에서 C/C++을 사용합니다. 개발자들은 언어 확장을 통해 소스 코드를 GPU 상에서 손쉽게 실행할 수 있게 됩니다.\n",
    "\n",
    "아래는 `.cu` 파일을 보여 줍니다. `.cu`는 CUDA 가속화 프로그램을 위한 파일 확장자입니다. 여기에서 두 개의 함수가 있는데 첫 번째는 CPU 상에서 실행되고 두 번째는 GPU 상에서 실행될 것입니다. 각각의 함수가 서로 어떻게 다른지, 정의되는 방식과 호출되는 방식에 유념하여 잘 살펴 보세요.\n",
    "\n",
    "```cpp\n",
    "void CPUFunction()\n",
    "{\n",
    "  printf(\"This function is defined to run on the CPU.\\n\");\n",
    "}\n",
    "\n",
    "__global__ void GPUFunction()\n",
    "{\n",
    "  printf(\"This function is defined to run on the GPU.\\n\");\n",
    "}\n",
    "\n",
    "int main()\n",
    "{\n",
    "  CPUFunction();\n",
    "\n",
    "  GPUFunction<<<1, 1>>>();\n",
    "  cudaDeviceSynchronize();\n",
    "}\n",
    "```\n",
    "\n",
    "아래에 코드의 중요 부분과 함께 가속화 컴퓨팅에서 사용하는 주요 용어들을 정리했습니다:\n",
    "\n",
    "`__global__ void GPUFunction()`\n",
    "* `__global__` 키워드는 해당 함수가 GPU에서 실행되며 **전역적**으로 호출될 수 있음을 의미합니다. 지금 문맥에서 전역적이란 CPU나 GPU, 그 누구에 의해서든 호출될 수 있다는 뜻입니다.\n",
    "* CPU에서 수행되는 코드는 종종 **호스트(host)** 코드라고 불립니다. GPU에서 수행되는 코드는 **디바이스(device)** 코드라고 합니다.\n",
    "* 리턴 타입은 `void`입니다. `__global__` 키워드로 정의된 함수는 반드시 `void` 타입을 리턴해야 합니다.\n",
    "\n",
    "`GPUFunction<<<1, 1>>>();`\n",
    "* GPU에서 실행되는 함수를 호출할 때, 통상적으로 이 함수를 **커널(kernel)**이라고 부르며 호출 대신에 **구동(launch)**이라는 표현을 씁니다. \n",
    "* 커널을 구동할 때에는 **실행 설정(execution configuration)**을 해주어야 합니다. 이것은 커널에 인수를 전달하는 부분 직전에서 `<<<...>>>` 꼴의 문법을 사용하여 이루어집니다. 실행 설정은 본 강좌에서 자세히 살펴 볼 것입니다. 지금은 일단 커널이 (`<<<1, 1>>>`의 첫 번째 값 1로부터 온) 1 개의 스레드 블록을 사용하고, 스레드 블록당 (`<<<1, 1>>>`의 두 번째 값 1로부터 온) 1 개의 스레드를 사용한다고만 이해하면 됩니다.\n",
    "\n",
    "`cudaDeviceSynchronize();`\n",
    "* 대부분의 C/C++ 코드와는 달리 커널 구동은 **비동기적(asynchronous)**입니다. 즉, CPU 코드는 커널 구동이 완료되는 것을 기다리지 않고 다음 코드를 실행합니다.\n",
    "* CUDA 런타임에서 제공되는 `cudaDeviceSynchronize`의 호출은 호스트(CPU) 코드로 하여금 디바이스(GPU) 코드의 완료를 기다리게 하며, 디바이스 코드가 완료되었을 때 비로소 CPU가 실행을 재개합니다.\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "qsp6SrRtEtyM"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습: Hello GPU 커널 작성\n",
    "\n",
    "[`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/01-hello/01-hello-gpu.cu) (<---- 소스 파일을 새로운 탭에서 열고 수정하려면 클릭하세요.)는 실제 동작하는 코드입니다. 이것은 두 개의 함수를 가지고 있고 둘 다 \"Hello from the CPU\"라는 메시지를 출력합니다. 여러분의 목표는 소스 파일에서 `helloGPU` 함수를 수정하여 GPU에서 실행되게 하고, 그 사실을 메시지로 출력하도록 만드는 것입니다. \n",
    "\n",
    "* 애플리케이션을 수정한 뒤, 아래와 같이 셀에서 컴파일하고 실행하세요. 코드 실행 셀에서 `CTRL+ENTER`를 누르면 코드가 실행된다는 점을 상기하세요. [`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/01-hello/01-hello-gpu.cu)의 주석을 참고하세요. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/01-hello/solutions/01-hello-gpu-solution.cu)을 참고하세요. 아래 명령을 이용해서 컴파일과 실행을 하기 전에 여러분이 수정한 파일을 저장해야 한다는 것을 잊지 마세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "AQuufZs9-2pf"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o hello-gpu 01-hello/01-hello-gpu.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "3dQcmrlk-2pi"
   },
   "source": [
    "[`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/01-hello/01-hello-gpu.cu)를 성공적으로 수정했다면 아래에 기술된 바와 같이 차례대로 수정하세요. 단, 하나를 수정할 때마다 위 셀의 명령을 (`CTRL+ENTER`을 눌러서) 실행해서 컴파일과 실행을 해보세요. 오류가 발생하면 오류 메시지를 주의 깊게 읽어 보세요. 오류 메시지에 익숙해지면 나중에 여러분의 가속화 코드를 작성할 때 큰 도움이 됩니다.\n",
    "\n",
    "* 커널 정의 부분의 `__global__` 키워드를 삭제하세요. 오류가 발생한 라인 번호를 기억하세요. 오류 메시지에 나온 \"configured\"의 의미가 무엇일까요? 끝나면 `__global__`을 원래대로 복구하세요.\n",
    "* 실행 설정을 삭제하세요. 여러분이 앞에서 이해한 \"configured\"의 의미가 여전히 유효한가요? 끝나면 실행 설정을 원래대로 복구하세요.\n",
    "* `cudaDeviceSynchronize` 호출을 삭제하세요. 컴파일과 실행에 앞서 무슨 일이 일어날 지 추측해 보세요. 커널은 비동기적으로 구동되고 `cudaDeviceSynchronize`는 호스트의 실행을 멈추고 커널 실행이 끝나기를 기다린다는 점을 상기하세요. 끝나면 `cudaDeviceSynchronize`을 원래대로 복구하세요.\n",
    "* `Hello from the GPU`가 `Hello from the CPU` 보다 **먼저** 출력되도록 `01-hello-gpu.cu`를 수정하세요.\n",
    "* `Hello from the GPU`가 **두 번** 출력되게 하되, 하나는 `Hello from the CPU` **이전**에, 다른 하나는 **이후**에 출력되도록 `01-hello-gpu.cu`를 수정하세요.\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "johsCXHIEJaA"
   },
   "source": [
    "---\n",
    "\n",
    "## 가속화 CUDA 코드 컴파일하고 실행하기\n",
    "\n",
    "본 섹션은 앞에서 `.cu` 프로그램을 컴파일하고 실행하기 위해 사용한 `nvcc` 명령어에 대하여 자세히 살펴 봅니다.\n",
    "\n",
    "CUDA 플랫폼은 [NVIDIA CUDA 컴파일러](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html)인 `nvcc`를 포함하여 배포됩니다. `nvcc`는 호스트 코드와 디바이스 코드를 포함한 CUDA 가속화 애플리케이션을 컴파일할 수 있습니다. 본 강좌의 목적상 `nvcc`에 대한 논의는 강좌에 필요한 정도로만 제한할 예정입니다. 본 강좌를 수료한 뒤, `nvcc`를 더 깊이 이해하고자 한다면 http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html 문서부터 읽어가시면 됩니다.\n",
    "\n",
    "`gcc` 사용자들에게 `nvcc`는 꽤 친숙하게 보일 것입니다. 예를 들어, `some-CUDA.cu`를 컴파일하려면 간단히 `nvcc -arch=sm_70 -o out some-CUDA.cu -run`이라고 입력하면 됩니다.\n",
    "\n",
    "* `nvcc`는 `nvcc` 컴파일러를 사용하기 위한 커맨드라인 명령어입니다.\n",
    "* `some-CUDA.cu`는 컴파일할 파일로서 전달됩니다.\n",
    "* `o` 플래그는 컴파일된 프로그램의 출력 파일 이름을 지정합니다.\n",
    "* `arch` 플래그는 파일이 어떤 **아키텍쳐**를 위해 컴파일되는가를 지정합니다. 여기에서 `sm_70`은 강좌에서 사용하는 볼타(Volta) GPU를 의미합니다. 보다 자세한 사항은 다음 문서를 참고하세요.\n",
    "    * [arch 플래그](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#options-for-steering-gpu-code-generation)\n",
    "    * [가상 아키텍처 기능](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#virtual-architectures)     \n",
    "    * [GPU 기능](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#gpu-feature-list) \n",
    "* 편의상 `run` 플래그를 추가하여 컴파일된 바이너리를 자동으로 실행하게 합니다.\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "vbLCloA1EK1g"
   },
   "source": [
    "\n",
    "---\n",
    "\n",
    "## CUDA 스레드 계층\n",
    "\n",
    "다음 슬라이드는 앞으로 공부할 내용을 개략적인 수준에서 시각적으로 보여줍니다. 주제를 보다 상세히 다루기에 앞서 슬라이드를 클릭하시면서 살펴 보세요. "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {
     "base_uri": "https://localhost:8080/",
     "height": 591
    },
    "colab_type": "code",
    "executionInfo": {
     "elapsed": 850,
     "status": "ok",
     "timestamp": 1541021739192,
     "user": {
      "displayName": "Hyungon Ryu KR",
      "photoUrl": "https://lh5.googleusercontent.com/-wtRHVbivPO4/AAAAAAAAAAI/AAAAAAAAABk/eM5PPjT2Row/s64/photo.jpg",
      "userId": "08585482373097815411"
     },
     "user_tz": -540
    },
    "id": "h3eOVG4C-2pk",
    "outputId": "628c2274-d4d4-4a61-b941-a8114676dfdc"
   },
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/e/2PACX-1vQYti_rVyWNNOccK6Slxd1VqazuqO5IhP17tmk-yTZAQfPEVpF14aZF9Vo3XkrDbFetNLTm_Pnk7JvD/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "TE46ymMk-2pm"
   },
   "source": [
    "## 병렬 커널 구동하기\n",
    "\n",
    "프로그래머는 실행 설정을 통해 다수의 GPU **스레드**에서 병렬로 실행될 커널의 세부사항을 지정합니다. 정확히 말하면 실행 설정은 얼마나 많은 스레드 그룹(**스레드 블록** 또는 **블록**이라고 합니다)이 있으며 각각의 스레드 그룹에는 얼마나 많은 스레드가 있는가를 나타냅니다. 문법은 다음과 같습니다.\n",
    "\n",
    "    <<<블록의 수, 블록당 스레드 수>>>\n",
    "\n",
    "**커널 코드는 구동될 때 설정된 모든 스레드 블록의 모든 스레드에서 동시에 실행됩니다.**\n",
    "\n",
    "따라서 `someKernel`이라는 커널이 있다고 할 때, 다음에 기술된 내용은 모두 참입니다.\n",
    "\n",
    "* `someKernel<<<1, 1>>>()`는 1 개의 스레드를 가진 스레드 블록 1 개에서 실행되는 설정이다. 따라서 커널은 1 회 수행된다.\n",
    "* `someKernel<<<1, 10>>>()`는 10 개의 스레드를 가진 스레드 블록 1 개에서 실행되는 설정이다. 따라서 커널은 10 회 수행된다.\n",
    "* `someKernel<<<10, 1>>>()`는 1 개의 스레드를 가진 스레드 블록 10 개에서 실행되는 설정이다. 따라서 커널은 10 회 수행된다.\n",
    "* `someKernel<<<10, 10>>>()`는 10 개의 스레드를 가진 스레드 블록 10 개에서 실행되는 설정이다. 따라서 커널은 100 회 수행된다.\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "tcIHvFFaENY3"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습: 병렬 커널 구동하기\n",
    "\n",
    "[`01-first-parallel.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/02-first-parallel/01-basic-parallel.cu)는 `This should be running in parallel`이라는 단순한 메시지를 출력합니다. 아래에 정리된 단계에 따라 코드를 수정하되, 먼저 코드가 GPU에서 실행되게 하고, 단일 스레드 블록에서 병렬로 실행되게 한 후, 마지막으로 다수의 블록에서도 병렬로 실행되도록 만드세요. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/02-first-parallel/solutions/01-basic-parallel-solution.cu)을 참고하세요. \n",
    "\n",
    "* `firstParallel` 함수를 수정하여 GPU 상에서 수행되는 CUDA 커널로서 구동하세요. 아래 셀의 명령어를 이용하여 `01-basic-parallel.cu`을 컴파일하고 실행하면 여전히 같은 메시지가 출력되어야 합니다. \n",
    "* `firstParallel` 커널을 수정하여 하나의 블록 안에 있는 5 개의 스레드에서 병렬로 실행되게 하세요. 컴파일하고 실행하면 메시지가 5 번 출력되어야 합니다.\n",
    "* `firstParallel` 커널을 다시 수정하여 이번에는 블록당 5 개의 스레드를 가진 블록 5 개에서 병렬로 실행되게 하세요. 컴파일하고 실행하면 메시지가 25 번 출력되어야 합니다."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "UCURaQ3K-2pp"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o basic-parallel 02-first-parallel/01-basic-parallel.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "T5_IuQ5s-2pt"
   },
   "source": [
    "---\n",
    "\n",
    "## CUDA에서 지원하는 스레드 계층 변수\n",
    "\n",
    "다음 슬라이드는 앞으로 공부할 내용을 개략적인 수준에서 시각적으로 보여줍니다. 주제를 보다 상세히 다루기에 앞서 슬라이드를 클릭하시면서 살펴 보세요. "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {
     "base_uri": "https://localhost:8080/",
     "height": 591
    },
    "colab_type": "code",
    "executionInfo": {
     "elapsed": 872,
     "status": "ok",
     "timestamp": 1541021747326,
     "user": {
      "displayName": "Hyungon Ryu KR",
      "photoUrl": "https://lh5.googleusercontent.com/-wtRHVbivPO4/AAAAAAAAAAI/AAAAAAAAABk/eM5PPjT2Row/s64/photo.jpg",
      "userId": "08585482373097815411"
     },
     "user_tz": -540
    },
    "id": "dNhAUx23-2pw",
    "outputId": "7af0ed1c-1b32-4cb1-dc90-c01160dcbc34"
   },
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/e/2PACX-1vSVS21bI-cje3Cqtxke-LHcvxk1ZxvZF-F35bgHSKfvNsvkGklCeqwlXHCDPJey5meZ1vTVYMqiF0UV/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "tEbiyQZV-2p1"
   },
   "source": [
    "---\n",
    "\n",
    "## 스레드 인덱스와 블록 인덱스 (Thread and Block Indices)\n",
    "\n",
    "스레드 블록 안에 있는 각 스레드에는 0부터 시작하는 인덱스가 할당됩니다. 또한, 각 블록 역시 0부터 시작하는 인덱스를 가지고 있습니다. 스레드가 모여 스레드 블록이 구성되듯이 블록이 모여 **그리드(grid)**가 됩니다. 그리드는 CUDA 스레드 계층 구조에서 최상위 개체에 해당합니다. 요약하면, CUDA 커널은 하나 이상의 블록으로 이루어진 그리드에서 실행되고, 각각의 블록은 1 개 이상의 스레드로 구성되는데 블록당 스레드의 갯수는 동일합니다.\n",
    "\n",
    "CUDA 커널은 (블록 안에서의) 스레드 인덱스를 나타내는 변수와 (그리드 안에서의) 블록 인덱스를 나타내는 변수를 가지고 있습니다. 이들을 각각 `threadIdx.x`와 `blockIdx.x`라고 합니다.\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "Ky-422iOEPyC"
   },
   "source": [
    "---\n",
    "\n",
    "\n",
    "### 실습:  특정한 스레드 및 블록 인덱스 사용하기\n",
    "\n",
    "[`01-thread-and-block-idx.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/03-indices/01-thread-and-block-idx.cu) 파일은 실패 메시지를 출력하는 커널을 포함하고 있습니다. 파일을 열고, 실행 설정을 어떻게 바꾸어야 성공 메시지가 출력될 것인지 살펴 보세요. 코드를 수정한 후 아래 셀을 이용하여 컴파일 및 실행을 하여 수정이 잘 되었는지 확인하세요. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/03-indices/solutions/01-thread-and-block-idx-solution.cu)을 참고하세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "wtYYaBsC-2p3"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o thread-and-block-idx 03-indices/01-thread-and-block-idx.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "dnsT8r2A-2p6"
   },
   "source": [
    "---\n",
    "\n",
    "## For 반복문 가속하기 (Accelerating For Loops)\n",
    "\n",
    "CPU-only 애플리케이션의 for 반복문은 가속화하기에 매우 적합한 대상입니다. 반복문으로 각 반복을 순차적으로 실행하는 대신에, 각 반복을 고유의 스레드에서 병렬로 실행되도록 하면 됩니다. 아래 반복문을 보면 반복이 몇 번 실행되는지와 각 반복에서 어떤 일이 일어나는가를 쉽게 알 수 있습니다.  \n",
    "\n",
    "```cpp\n",
    "int N = 2<<20;\n",
    "for (int i = 0; i < N; ++i)\n",
    "{\n",
    "  printf(\"%d\\n\", i);\n",
    "}\n",
    "```\n",
    "\n",
    "이 반복문을 병렬화하기 위해서 두 단계의 과정이 필요합니다.\n",
    "\n",
    "* 커널은 반복 중 **1 회**만 수행하도록 작성합니다.\n",
    "* 커널은 다른 커널의 동작 여부를 알 수 없으므로 정확한 횟수로 커널이 수행되도록 실행 설정을 해야 합니다. 즉, 반복문의 반복 횟수만큼 커널이 수행되어야 합니다.\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "jk0RRDb4ER0T"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습: 단일 스레드 블록으로 for 반복문 가속하기\n",
    "\n",
    "[`01-single-block-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/04-loops/01-single-block-loop.cu) 내부의 `loop` 함수는 0 부터 9 까지의 숫자를 순차적으로 출력하는 반복문을 수행합니다. `loop` 함수를 수정해서 CUDA 커널로 만들고 N 번의 반복을 병렬로 실행하도록 구동하세요. 성공적으로 수정된 경우 여전히 0 부터 9 까지의 숫자가 출력되어야 합니다. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/04-loops/solutions/01-single-block-loop-solution.cu)을 참고하세요. "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "5UTqfHyY-2p8"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o single-block-loop 04-loops/01-single-block-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "rBqALzTV-2p_"
   },
   "source": [
    "---\n",
    "\n",
    "## 병렬 스레드 협력 (Coordinating Parallel Threads)\n",
    "\n",
    "다음 슬라이드는 앞으로 공부할 내용을 개략적인 수준에서 시각적으로 보여줍니다. 주제를 보다 상세히 다루기에 앞서 슬라이드를 클릭하시면서 살펴 보세요. \n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {
     "base_uri": "https://localhost:8080/",
     "height": 591
    },
    "colab_type": "code",
    "executionInfo": {
     "elapsed": 860,
     "status": "ok",
     "timestamp": 1541021753545,
     "user": {
      "displayName": "Hyungon Ryu KR",
      "photoUrl": "https://lh5.googleusercontent.com/-wtRHVbivPO4/AAAAAAAAAAI/AAAAAAAAABk/eM5PPjT2Row/s64/photo.jpg",
      "userId": "08585482373097815411"
     },
     "user_tz": -540
    },
    "id": "qjrwDdip-2qA",
    "outputId": "57bbf6e3-8784-4fff-91fb-ae861d15ee7a"
   },
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/e/2PACX-1vSfi8LAinJ1RqTzlB2vRsAcDzCCk9gZov5rQODN5rtRMPt57UizCVv5LSVZ5WLxGtrsMm7FIkLb0wMR/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "Ts07pSKp-2qD"
   },
   "source": [
    "---\n",
    "\n",
    "## 병렬화 수준을 높이기 위해 블록 차원 사용하기 (Block Dimensions)\n",
    "\n",
    "하나의 스레드 블록에는 최대 1024 개까지의 스레드만 들어갈 수 있습니다. 가속화 애플리케이션의 병렬화 수준을 높이기 위해서는 다수의 스레드 블록이 서로 협력하도록 해야 합니다.\n",
    "\n",
    "CUDA 커널은 하나의 블록에 들어있는 스레드의 수를 저장하는 `blockDim.x`라는 특별한 변수를 가지고 있습니다. 이 변수를 `blockIdx.x` 및 `threadIdx.x`와 함께 사용하여 블록과 스레드의 병렬 수행을 조직화함으로써 병렬화 수준을 높일 수 있는데, 이때 `threadIdx.x + blockIdx.x * blockDim.x`와 같은 수식을 사용합니다. 상세한 예는 다음과 같습니다. (역주: 이는 처리할 데이터를 1 차원으로 보았을 때의 예입니다. 2 차원의 영상 데이터 등을 처리한다고 할 때에는 `threadIdx.x + blockIdx.x * blockDim.x` 뿐 아니라 `threadIdx.y + blockIdx.y * blockDim.y`도 함께 사용됩니다.) \n",
    "\n",
    "실행 설정 `<<<10, 10>>>`은 블록 당 10 개의 스레드(역주: blockDim.x는 10)를 가진 블록 10 개로 이루어진, 총 100 개의 스레드로 구성된 그리드를 구동합니다. 따라서 0부터 99까지의 인덱스를 가진 값들 중에서 고유한 하나를 스레드에게 주어 계산시킬 수 있을 것입니다. (역주: 즉, 서로 다른 스레드가 같은 인덱스의 값을 계산하면 중복이 되므로 중복이 되지 않도록 각 스레드에게 일감을 배분한다는 뜻입니다.) \n",
    "\n",
    "* 블록의 `blockIdx.x`가 0이면 `blockIdx.x * blockDim.x`는 0이 됩니다. 이 0을 가능한 `threadIdx.x` 값인 0부터 9의 값에 더하면 100 개의 스레드 그리드 중 0부터 9까지의 인덱스를 얻을 수 있습니다.\n",
    "* 블록의 `blockIdx.x`가 1이면 `blockIdx.x * blockDim.x`는 10이 됩니다. 이 10을 가능한 `threadIdx.x` 값인 0부터 9의 값에 더하면 100 개의 스레드 그리드 중 10부터 19까지의 인덱스를 얻을 수 있습니다.\n",
    "* 블록의 `blockIdx.x`가 5이면 `blockIdx.x * blockDim.x`는 50이 됩니다. 이 50을 가능한 `threadIdx.x` 값인 0부터 9의 값에 더하면 100 개의 스레드 그리드 중 50부터 59까지의 인덱스를 얻을 수 있습니다.\n",
    "* 블록의 `blockIdx.x`가 9이면 `blockIdx.x * blockDim.x`는 90이 됩니다. 이 90을 가능한 `threadIdx.x` 값인 0부터 9의 값에 더하면 100 개의 스레드 그리드 중 90부터 99까지의 인덱스를 얻을 수 있습니다.\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "iqwjsC4VEUk8"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습:  다중 스레드 블록에서 for 반복문 가속화하기\n",
    "\n",
    "[`02-multi-block-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/04-loops/02-multi-block-loop.cu) 내부의 `loop` 함수는 0 부터 9 까지의 숫자를 순차적으로 출력하는 for 반복문을 수행합니다. `loop` 함수를 수정해서 CUDA 커널로 만들고 `N` 번의 반복을 병렬로 실행하도록 구동하세요. 성공적으로 수정된 경우 여전히 0 부터 9 까지의 숫자가 출력되어야 합니다. 이 실습을 위한 추가 제약사항으로서, 최소 2 개 이상의 스레드 블록을 구동하도록 실행 설정을 하세요. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/04-loops/solutions/02-multi-block-loop-solution.cu)을 참고하세요. \n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "9mzId9hB-2qD"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o multi-block-loop 04-loops/02-multi-block-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "IUpYqTHE-2qI"
   },
   "source": [
    "---\n",
    "\n",
    "## GPU와 CPU가 사용하는 메모리 할당하기 (Allocating Memory)\n",
    "\n",
    "버전 6 이상의 최근 버전 CUDA에서는 CPU 호스트와 GPU 디바이스가 사용하는 메모리를 보다 쉽게 할당할 수 있습니다. 가속화 애플리케이션의 최적 성능을 위한 다양한 [중급 및 고급 메모리 관리 기법들](http://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#memory-optimizations) 이 있기는 하지만, 지금 소개하는 기본 CUDA 메모리 관리 기법만으로도 별다른 어려움 없이 CPU-only 애플리케이션에 비해 훨씬 훌륭한 성능을 얻을 수 있습니다.\n",
    "\n",
    "메모리를 할당하고 해제하며, 호스트와 디바이스 코드가 참조할 수 있는 포인터를 얻기 위해서는 아래 예제와 같이 `malloc`과 `free`를 `cudaMallocManaged`와 `cudaFree`로 바꾸어 주기만 하면 됩니다.\n",
    "\n",
    "```cpp\n",
    "// CPU-only\n",
    "\n",
    "int N = 2<<20;\n",
    "size_t size = N * sizeof(int);\n",
    "\n",
    "int *a;\n",
    "a = (int *)malloc(size);\n",
    "\n",
    "// Use `a` in CPU-only program.\n",
    "\n",
    "free(a);\n",
    "```\n",
    "\n",
    "```cpp\n",
    "// Accelerated\n",
    "\n",
    "int N = 2<<20;\n",
    "size_t size = N * sizeof(int);\n",
    "\n",
    "int *a;\n",
    "// Note the address of `a` is passed as first argument.\n",
    "cudaMallocManaged(&a, size);\n",
    "\n",
    "// Use `a` on the CPU and/or on any GPU in the accelerated system.\n",
    "\n",
    "cudaFree(a);\n",
    "```\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "azS9DkRREXQc"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습: 호스트와 디바이스에서 배열 다루기\n",
    "\n",
    "[`01-double-elements.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/05-allocate/01-double-elements.cu) 프로그램은 배열을 할당한 후 호스트 상에서 정수값으로 초기화하고 각각의 값을 GPU 상에서 2로 곱합니다. 그리고 다시 호스트 상에서 이배수 연산의 성공 여부를 확인합니다. 현재 이 프로그램은 잘 작동하지 않습니다. 포인터 `a`를 이용하여 호스트와 디바이스 양쪽에서 모두 메모리 접근이 가능해야 하지만 배열이 `malloc`으로 할당되어 있어서 호스트에서만 접근 가능합니다. 애플리케이션을 다음 조건에 따라 수정하세요. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/05-allocate/solutions/01-double-elements-solution.cu)을 참고하세요.           \n",
    "\n",
    "* `a`는 호스트 코드와 디바이스 코드 양쪽에서 모두 사용할 수 있어야 합니다.\n",
    "* `a`가 참조하는 메모리는 올바르게 해제되어야 합니다.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "VkW3l2KH-2qK"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o double-elements 05-allocate/01-double-elements.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "-HuIsBt7-2qQ"
   },
   "source": [
    "## 그리드 크기와 작업량 간의 불일치\n",
    "\n",
    "다음 슬라이드는 앞으로 공부할 내용을 개략적인 수준에서 시각적으로 보여줍니다. 주제를 보다 상세히 다루기에 앞서 슬라이드를 클릭하시면서 살펴 보세요. "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {
     "base_uri": "https://localhost:8080/",
     "height": 591
    },
    "colab_type": "code",
    "executionInfo": {
     "elapsed": 863,
     "status": "ok",
     "timestamp": 1541021763717,
     "user": {
      "displayName": "Hyungon Ryu KR",
      "photoUrl": "https://lh5.googleusercontent.com/-wtRHVbivPO4/AAAAAAAAAAI/AAAAAAAAABk/eM5PPjT2Row/s64/photo.jpg",
      "userId": "08585482373097815411"
     },
     "user_tz": -540
    },
    "id": "mhHEsYU3-2qQ",
    "outputId": "87d4136e-1471-46bf-b358-e840f5d30bcb"
   },
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/e/2PACX-1vTn2HX1FyUO94g5TxgBm0C7pu-_5UXPwYtMkhGLnqgs0-2Y1g8CE3YCuZuob25wrrXz0x8cT9_XxyBl/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "Jdwoekk0-2qT"
   },
   "source": [
    "---\n",
    "## 필요한 스레드 수와 블록 설정 간의 불일치 해결하기 (Grid Size Work Amount Mismatch)\n",
    "\n",
    "때로는 반복문을 병렬화하기 위해 필요한 정확한 스레드 갯수를 실행 설정 과정에서 지정할 수 없는 경우가 있을 수 있습니다.\n",
    "\n",
    "이러한 사례는 최적 블록 크기를 결정할 때 흔히 볼 수 있습니다. 예를 들어, GPU 하드웨어의 특성 때문에 블록에 들어가는 스레드의 수를 32의 배수로 설정해야만 좋은 성능이 나오는 경우가 있습니다. 32의 배수인 256 개의 스레드를 가지는 블록을 여러 개 구동하고 싶고, 1000 개의 병렬 작업을 실행해야 하는 경우를 생각해 봅시다. 이 경우 정확히 1000 개의 스레드를 가지는 블록 수를 정할 수 없습니다. 32의 배수로는 정확히 1000을 만들 수 없기 때문입니다.\n",
    "\n",
    "이 시나리오는 다음과 같은 방법으로 해결할 수 있습니다.\n",
    "\n",
    "* 할당된 작업을 수행하는 데 필요한 스레드 수보다 **더 많은** 수의 스레드를 만들도록 실행 설정을 작성합니다.\n",
    "* 처리해야 하는 데이터의 총 크기 또는 작업을 위해 필요한 총 스레드 수를 `N`이라고 하고 이 값을 커널에 인수로 넘겨줍니다.\n",
    "* `tid+bid*bdim`를 이용하여 스레드의 인덱스를 계산한 후, 이 인덱스 값이 `N`을 초과하지 않는지 확인하여 초과하지 않는 경우에만 커널의 관련 작업을 수행합니다.\n",
    "\n",
    "`N` 값과 블록 당 스레드 갯수를 알고 있고, 그리드의 스레드 갯수와 `N`이 정확히 일치하는 것을 보장할 수 없을 때의 통상적 예제 코드는 아래와 같습니다. 이 코드는 최소 `N` 개 이상의 스레드가 존재하되, 최대 1 개의 추가적 블록에 해당하는 스레드까지만 만들도록 되어 있습니다.\n",
    "\n",
    "```cpp\n",
    "// Assume `N` is known\n",
    "int N = 100000;\n",
    "\n",
    "// Assume we have a desire to set `threads_per_block` exactly to `256`\n",
    "size_t threads_per_block = 256;\n",
    "\n",
    "// Ensure there are at least `N` threads in the grid, but only 1 block's worth extra\n",
    "size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;\n",
    "\n",
    "some_kernel<<<number_of_blocks, threads_per_block>>>(N);\n",
    "```\n",
    "\n",
    "위 코드의 실행 설정은 `N` 보다 큰 갯수의 스레드를 가지는 그리드를 만들기 때문에 \"추가된\" 스레드 중 하나가 실행될 때 인덱스 범위를 벗어나는 데이터 접근을 시도하지 않도록 주의를 기울여야 합니다. \n",
    "\n",
    "\n",
    "```cpp\n",
    "__global__ some_kernel(int N)\n",
    "{\n",
    "  int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
    "\n",
    "  if (idx < N) // Check to make sure `idx` maps to some value within `N`\n",
    "  {\n",
    "    // Only do work if it does\n",
    "  }\n",
    "}\n",
    "```\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "9STXkAsmEZuR"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습: 실행 설정 불일치일 때 for 반복문 가속화하기\n",
    "\n",
    "[`02-mismatched-config-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/05- allocate/02-mismatched-config-loop.cu) 프로그램은 `cudaMallocManaged`를 이용하여 1000 개의 정수를 원소로하는 배열을 메모리에 할당합니다. 그리고 배열의 모든 원소를 CUDA 커널에서 초기화합니다. `N`과 `threads_per_block`은 알고 있다고 가정합니다. 여러분의 임무는 다음에 있는 두 가지 목표를 달성하는 것입니다. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/05-allocate/solutions/02-mismatched-config-loop-solution.cu)을 참고하세요.           \n",
    "\n",
    "* 최소한 `a`의 원소의 수만큼 스레드의 수를 보장하는 값을 `number_of_blocks`에 대입하세요.\n",
    "* `initializeElementsTo` 커널을 수정하여, 인덱스 범위를 벗어나는 원소에 대하여 작업을 시도하는 일이 없도록 만드세요.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "YPM-vM85-2qU"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o mismatched-config-loop 05-allocate/02-mismatched-config-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "w-L7nCSq-2qY"
   },
   "source": [
    "---\n",
    "\n",
    "## 그리드 폭 반복문 - 그리드 폭만큼 건너뛰며(stride) 반복하기 \n",
    "\n",
    "다음 슬라이드는 앞으로 공부할 내용을 개략적인 수준에서 시각적으로 보여줍니다. 주제를 보다 상세히 다루기에 앞서 슬라이드를 클릭하시면서 살펴 보세요. "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {
     "base_uri": "https://localhost:8080/",
     "height": 591
    },
    "colab_type": "code",
    "executionInfo": {
     "elapsed": 1101,
     "status": "ok",
     "timestamp": 1541021774215,
     "user": {
      "displayName": "Hyungon Ryu KR",
      "photoUrl": "https://lh5.googleusercontent.com/-wtRHVbivPO4/AAAAAAAAAAI/AAAAAAAAABk/eM5PPjT2Row/s64/photo.jpg",
      "userId": "08585482373097815411"
     },
     "user_tz": -540
    },
    "id": "Y5Gngsnw-2qZ",
    "outputId": "ecc2390d-f8fa-413e-9c92-044dc3fb0cf7"
   },
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/e/2PACX-1vTSfcPagyv7ObRnhygFnKrvDIDa-wUuc3yR-qs7xd4gQxProMOqXzNqe8y9vz711cLIbPp1qYJc7R3l/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "QCgGJgJk-2qd"
   },
   "source": [
    "---\n",
    "\n",
    "## 그리드보다 큰 데이터 집합\n",
    "\n",
    "최고 성능의 실행 설정을 위한 선택에 의해서이든지 아니면 그저 필요에 의해서든지, 그리드의 스레드 갯수가 데이터 집합의 크기보다 작은 경우가 있기 마련입니다. 1000 개의 원소를 가진 배열과 250 개의 스레드를 가진 그리드를 생각해 봅시다. 여기에서 그리드의 각 스레드는 4 번씩 사용되어야 합니다. 이를 달성하기 위한 통상적인 방법은 커널에서 **그리드 폭 반복문(grid-stride loop)**을 사용하는 것입니다. \n",
    "\n",
    "그리드 폭 반복문에서 각 스레드는 `tid+bid*bdim`를 이용하여 고유의 인덱스를 계산합니다. 그리고는 방금 계산한 인덱스 값에 그리드의 스레드 갯수를 더하여 다음 계산을 합니다. 이 과정을 배열의 범위를 벗어나기 전까지 반복합니다. 예를 들어 500 개의 원소를 가진 배열과 250 개의 스레드를 가진 그리드가 있을 때, 인덱스 20을 가진 스레드는 다음과 같은 과정을 밟을 것입니다.\n",
    "\n",
    "* 500 개의 원소를 가진 배열에서 20 번째 원소를 처리합니다.\n",
    "* 인덱스를 그리드의 폭인 250 만큼 건너뛰어 인덱스 270을 얻습니다.\n",
    "* 500 개의 원소를 가진 배열에서 270 번째 원소를 처리합니다.\n",
    "* 인덱스를 그리드의 폭인 250 만큼 건너뛰어 인덱스 520을 얻습니다.\n",
    "* 520은 인덱스 범위를 초과하므로 스레드는 작업을 멈춥니다.\n",
    "\n",
    "CUDA는 그리드 안에 있는 블록 수를 나타내는 변수 `gridDim.x`를 제공합니다. 그리드의 총 스레드 수는 단순히 그리드의 블록 수에 블록 당 스레드 수를 곱한 `gridDim.x * blockDim.x`로부터 얻을 수 있습니다. 이 내용을 기억하면서 아래 예제를 보세요.\n",
    "\n",
    "```cpp\n",
    "__global void kernel(int *a, int N)\n",
    "{\n",
    "  int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;\n",
    "  int gridStride = gridDim.x * blockDim.x;\n",
    "\n",
    "  for (int i = indexWithinTheGrid; i < N; i += gridStride)\n",
    "  {\n",
    "    // do work on a[i];\n",
    "  }\n",
    "}\n",
    "```\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "LCZAOe91Ecm-"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습: 그리드 폭 반복문을 이용하여 그리드보다 큰 배열 다루기\n",
    "\n",
    "[`03-grid-stride-double.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/05-allocate/03-grid- stride-double.cu) 프로그램을 수정하여 `doubleElements` 커널에서 그리드 폭 반복문을 사용하세요. 이렇게 함으로써 `N`보다 작은 그리드가 배열 전체를 처리하기 위해 스레드를 재사용할 수 있게 합니다. 프로그램은 배열의 모든 원소들이 두 배가 되었는지 아닌지를 출력할 것입니다. 지금은 `FALSE`를 출력할 것입니다. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/05-allocate/solutions/03-grid-stride-double-solution.cu)을 참고하세요.           "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "nxhIGVyl-2qg"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o grid-stride-double 05-allocate/03-grid-stride-double.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "6eIog7kp-2qm"
   },
   "source": [
    "---\n",
    "\n",
    "## 오류 처리하기  (Error Handling)\n",
    "\n",
    "다른 애플리케이션과 마찬가지로 CUDA 역시 오류 처리가 필수적입니다. 대부분이라고까지는 할 수 없지만 많은 CUDA 함수가 `cudaError_t` 타입의 값을 리턴합니다. ([메모리 관리 함수의 예](http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY)를 확인해 보세요.) 이 값을 이용해서 함수 호출 중에 오류가 발생했는지를 알 수 있습니다. 아래 코드는 `cudaMallocManaged`를 호출할 때 오류를 처리하는 예제입니다.\n",
    "\n",
    "```cpp\n",
    "cudaError_t err;\n",
    "err = cudaMallocManaged(&a, N)                    // Assume the existence of `a` and `N`.\n",
    "\n",
    "if (err != cudaSuccess)                           // `cudaSuccess` is provided by CUDA.\n",
    "{\n",
    "  printf(\"Error: %s\\n\", cudaGetErrorString(err)); // `cudaGetErrorString` is provided by CUDA.\n",
    "}\n",
    "```\n",
    "\n",
    "`void`를 리턴하도록 되어 있는 구동 커널은 `cudaError_t` 타입의 값을 리턴하지 않습니다. 실행 설정의 오류와 같이 커널 구동 단계에서 발생하는 오류를 점검하기 위해서 CUDA는 `cudaGetLastError` 함수를 제공하는데 이 함수는 `cudaError_t` 타입의 값을 리턴합니다.\n",
    "\n",
    "```cpp\n",
    "/*\n",
    " * This launch should cause an error, but the kernel itself\n",
    " * cannot return it.\n",
    " */\n",
    "\n",
    "someKernel<<<1, -1>>>();  // -1 is not a valid number of threads.\n",
    "\n",
    "cudaError_t err;\n",
    "err = cudaGetLastError(); // `cudaGetLastError` will return the error from above.\n",
    "if (err != cudaSuccess)\n",
    "{\n",
    "  printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
    "}\n",
    "```\n",
    "\n",
    "마지막으로, 비동기 커널의 실행 중에 나오는 오류와 같이 비동기적으로 발생하는 오류를 잡아내기 위해서는 뒤이어 실행되는 동기화를 위한 CUDA 런타임 API 호출에서 리턴된 상태값을 검사하는 것이 필수적입니다. 이러한 동기화 함수에는 앞에서 본 `cudaDeviceSynchronize`가 있는데 이 함수는 이전에 구동된 커널 중 실패한 것들의 오류를 리턴합니다. \n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "ZQoG8e_kEfgX"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습: 오류 처리 추가하기\n",
    "\n",
    "[`01-add-error-handling.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/06-errors/01-add-error- handling.cu) 프로그램을 컴파일하고 실행하면 배열의 원소를 이배수하는 작업이 성공하지 못했다는 결과를 출력합니다. 하지만 관련된 오류를 출력하지는 않습니다. CUDA 오류를 처리하도록 프로그램을 수정함으로써 어디가 잘못되었는지를 알 수 있게 하고 디버깅을 효과적으로 할 수 있도록 하세요. CUDA 함수 호출 중에 나오는 동기적 오류와 커널 실행 중에 나오는 비동기적 오류 모두를 검사해야 합니다. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/06-errors/solutions/01-add-error-handling-solution.cu)을 참고하세요.           \n",
    "\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "VkSeAzgM-2qm"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o add-error-handling 06-errors/01-add-error-handling.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "BQ7ECuIj-2qp"
   },
   "source": [
    "---\n",
    "\n",
    "### CUDA 오류 처리 함수\n",
    "\n",
    "오류 검사를 위하여 CUDA 함수를 래핑하는 매크로를 만들면 도움이 됩니다. 아래에 예제가 있으니 앞으로의 실습에서 자유로이 사용하세요.\n",
    "\n",
    "```cpp\n",
    "#include <stdio.h>\n",
    "#include <assert.h>\n",
    "\n",
    "inline cudaError_t checkCuda(cudaError_t result)\n",
    "{\n",
    "  if (result != cudaSuccess) {\n",
    "    fprintf(stderr, \"CUDA Runtime Error: %s\\n\", cudaGetErrorString(result));\n",
    "    assert(result == cudaSuccess);\n",
    "  }\n",
    "  return result;\n",
    "}\n",
    "\n",
    "int main()\n",
    "{\n",
    "\n",
    "/*\n",
    " * The macro can be wrapped around any function returning\n",
    " * a value of type `cudaError_t`.\n",
    " */\n",
    "\n",
    "  checkCuda( cudaDeviceSynchronize() )\n",
    "}\n",
    "```\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "jFnTynTs-TAv"
   },
   "source": [
    "\n",
    "---\n",
    "\n",
    "## 요약\n",
    "\n",
    "이제 여러분은 아래와 같은 학습 목표를 달성하셨습니다.\n",
    "\n",
    "* CPU 함수를 호출하고 GPU **커널(kernel)을 구동(launch)**하는 C/C++ 프로그램을 작성, 컴파일, 실행하기\n",
    "* **실행 설정(execution configuration)**을 통해 병렬 **스레드 계층(thread hierarchy)** 제어하기 \n",
    "* 직렬로 실행되는 반복문을 GPU 상에서 병렬로 실행되도록 수정하기\n",
    "* CPU와 GPU를 위한 메모리를 할당하고 해제하기\n",
    "* CUDA 코드에서 생성된 오류 처리하기\n",
    "\n",
    "이제 강좌의 마지막 목표를 다룰 차례입니다.\n",
    "\n",
    "* CPU-only 애플리케이션 가속화하기\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "lDG3-jZ_Eiyh"
   },
   "source": [
    "---\n",
    "\n",
    "### 마지막 실습: 벡터합 가속화하기\n",
    "\n",
    "이번 과제는 여러분이 지금까지 배운 내용을 활용할 기회가 될 것입니다. 비록 복잡한 프로그램은 아니지만 CPU-only 벡터합 프로그램을 가속화하는 내용을 포함하는데, CUDA를 이용한 GPU 가속 애플리케이션에 대하여 배운 것에 집중할 기회입니다. 실습을 마치고, 관심과 여유가 있다면 *고급 주제* 섹션에서 보다 복잡한 코드를 포함한 과제에 도전해 보세요.\n",
    "\n",
    "[`01-vector-add.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/07-vector-add/01-vector-add.cu) 프로그램은 제대로 동작하는 CPU-only 벡터합 계산을 포함하고 있습니다. `addVectorsInto` 함수가 GPU 상의 CUDA 커널로 동작하여 작업을 병렬로 처리하도록 코드를 수정하세요. 아래에 리스트로 명시된 작업이 수행되어야 합니다. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/07-vector-add/solutions/01-vector-add-solution.cu)을 참고하세요.\n",
    "\n",
    "* CUDA 커널이 되도록 `addVectorsInto`를 확장하세요.\n",
    "* `addVectorsInto`가 CUDA 커널로서 구동되도록 실행 설정을 하세요.\n",
    "* 세 개의 벡터 `a`, `b`, `result`가 호스트와 디바이스 코드에 의해 사용될 수 있도록 메모리 할당과 해제 부분을 수정하세요.\n",
    "* `addVectorsInto`의 내용을 수정하세요. `addVectorsInto`는 각각의 스레드에서 실행되어야 하고 각 스레드는 입력 벡터에 대하여 하나의 스레드에 해당하는 작업만을 수행해야 합니다. 스레드가 인덱스 범위를 초과하여 원소에 접근하지 않는지 확인하고 스레드가 입력 벡터의 원소를 하나 이상 처리해야 하는지도 잘 살펴 보세요.\n",
    "* CUDA 코드가 아무런 메시지 없이 실패할 수 있는 곳마다 오류 처리를 추가하세요."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "Hu9Rp_Wx-2qq"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o vector-add 07-vector-add/01-vector-add.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "Xf45qdtU-2qu"
   },
   "source": [
    "---\n",
    "\n",
    "## 고급 주제\n",
    "\n",
    "다음에 나오는 실습은 관심과 여유가 있는 분들을 위한 추가적인 도전 과제입니다. 고급 기법을 사용해야 하고 도움말은 적습니다. 여러분의 성장을 위한 어렵지만 멋진 과제들입니다.\n",
    "\n",
    "---\n",
    "\n",
    "## 2 차원 또는 3 차원 그리드와 블록\n",
    "\n",
    "그리드와 블록은 3 차원까지 정의할 수 있습니다. 다차원으로 정의한다고 해서 성능에 영향을 미치는 것은 아닙니다만, 2 차원 행렬과 같은 다차원 데이터를 다루는 데 편리합니다. 그리드나 블록을 다차원으로 정의하려면 CUDA에서 제공하는 `dim3` 타입을 사용합니다.\n",
    "\n",
    "```cpp\n",
    "dim3 threads_per_block(16, 16, 1);\n",
    "dim3 number_of_blocks(16, 16, 1);\n",
    "someKernel<<<number_of_blocks, threads_per_block>>>();\n",
    "```\n",
    "\n",
    "위와 같은 예에서는 `someKernel` 내부에서 `gridDim.x`, `gridDim.y`, `blockDim.x`, `blockDim.y` 변수값이 모두 16을 가질 것입니다.\n",
    "\n",
    "---\n",
    "\n",
    "\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "mRzM5CP4EmnG"
   },
   "source": [
    "### 실습: 2 차원 행렬곱 가속화하기\n",
    "\n",
    "[`01-matrix-multiply-2d.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/08-matrix-multiply/01-matrix-multiply-2d.cu) 프로그램은 완벽히 동작하는 호스트 함수인 `matrixMulCPU`를 포함하고 있습니다. 여러분이 할 일은 `matrixMulGPU` CUDA 커널을 만드는 일입니다. 소스 코드는 두 함수를 모두 이용하여 행렬곱을 수행해야 하며 결과를 서로 비교하여 여러분이 만든 CUDA 커널을 검사해야 합니다. 아래 리스트에 명시된 내용을 참고하세요. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/08-matrix-multiply/solutions/01-matrix-multiply-2d-solution.cu)을 참고하세요.\n",
    "\n",
    "* `x`, `y` 차원이 1 보다 큰 값을 가지는 `dim3` 타입의 값을 인수로 하여 실행 설정을 해야 합니다.\n",
    "* 커널의 본문에서는 기존과 같이 실행 스레드의 그리드 내 고유 인덱스 값을 결정해야 합니다. 단, 이제는 그리드의 x 축과 y 축에 해당하는 두 개의 인덱스를 사용해야 합니다.  "
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "PXyV80DO-2qv"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o matrix-multiply-2d 08-matrix-multiply/01-matrix-multiply-2d.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "v1Y4otxO-2qx"
   },
   "source": [
    "---\n",
    "\n",
    "### 실습: 열전도(heat conduction) 애플리케이션 가속화하기\n",
    "\n",
    "이번 실습에서 여러분은 은의 2 차원 공간 열전도를 시뮬레이션하는 애플리케이션을 가속화할 것입니다.\n",
    "\n",
    "[`01-heat-conduction.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/09-heat/01-heat-conduction.cu) 프로그램에 있는 `step_kernel_mod` 함수가 GPU에서 실행되도록 수정하고 `main` 함수에서는 CPU와 GPU가 사용할 데이터를 적절히 할당하세요. `step_kernel_ref`는 CPU에서 동작하며 오류 점검에 사용됩니다. 이 코드는 부동소수점 연산을 포함하고 있으므로 다른 프로세서에서 실행하거나 같은 프로세서에서의 단순한 순서바꿈에 의해서도 약간 다른 결과를 도출할 수 있습니다. 이러한 이유로 인하여 오류 점검 코드는 정확한 값의 일치 대신에 임계값을 사용합니다. 어떻게 해야할 지 모르거나 작업을 확인해야 한다면 [솔루션](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-kr/09-heat/solutions/01-heat-conduction-solution.cu)을 참고하세요.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "colab": {},
    "colab_type": "code",
    "id": "iD1XzfTc-2qy"
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o heat-conduction 09-heat/01-heat-conduction.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "JPk7_1F0-2q1"
   },
   "source": [
    "> 열전달(heat conduction) CPU 소스 코드의 원본 출전은 휴스턴 대학의 [An OpenACC Example Code for a C-based heat conduction code](http://docplayer.net/30411068-An-openacc-example-code-for-a-c-based-heat-conduction-code.html) 임을 밝혀 둡니다. "
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "colab_type": "text",
    "id": "sFjczZATD_Uc"
   },
   "source": [
    "![CUDA](./images/CUDA_Logo.jpg)"
   ]
  }
 ],
 "metadata": {
  "colab": {
   "collapsed_sections": [],
   "name": "01_AC_CUDA_C-kr.ipynb",
   "provenance": [],
   "version": "0.3.2"
  },
  "kernelspec": {
   "display_name": "Python 2",
   "language": "python",
   "name": "python2"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 2
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython2",
   "version": "2.7.12"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 1
}
